SWDEV-546808 optimize match_any() with dpp wave_rot1 #183
Conversation
| retval = 1; | ||
| for (int i = 1; i < static_cast<int>(warpSize); i++) { | ||
| if constexpr (sizeof(value) == 8) | ||
| dill_.ill = __builtin_amdgcn_mov_dpp(dill_.ill, 0x134, 0xf, 0xf, 0); //wave_rol1 |
There was a problem hiding this comment.
What does the comment "wave_rol1" mean? Also, the magic constants need comments like (/* foo= */ 0x134). Or even better would be const unsigned int meaningfulName = 0x134.
There was a problem hiding this comment.
It is the dpp_ctrl. I tried documenting here:
https://medium.com/@hashem.hashemi/amds-dpp-operations-for-super-fast-reduction-29f9aa888376
There was a problem hiding this comment.
But that's just an external personal log. Having meaningful names (with comments if necessary) at the top of this file will be much better. Could be const integers, or could be an enum class.
There was a problem hiding this comment.
More comments added in new check-in. Let me know if it makes it clearer.
| dill_.i[0] = __builtin_amdgcn_mov_dpp(dill_.i[0], 0x134, 0xf, 0xf, 0); //wave_rol1 | ||
| retval |= ((decltype(mask))(dill_.val == value)) << i; | ||
| } | ||
| int rotv = __lane_id(); |
There was a problem hiding this comment.
What does rotv mean? More comments will be useful for future reference, like what is the overall strategy here.
There was a problem hiding this comment.
Its the rotate amount. I added comment to code.
| retval = 1; | ||
| for (int i = 1; i < static_cast<int>(warpSize); i++) { | ||
| if constexpr (sizeof(value) == 8) | ||
| dill_.ill = __builtin_amdgcn_mov_dpp(dill_.ill, 0x134, 0xf, 0xf, 0); //wave_rol1 |
There was a problem hiding this comment.
I need to read up on this, but just to be sure, do these _dpp intrinsics work with both wave64 and wave32?
There was a problem hiding this comment.
I don't see why it wouldn't work for wave32. That is why I put the rotate amount as static_cast(warpSize).
| } | ||
| } | ||
| } else { | ||
| union dill { unsigned int i[2]; unsigned long long ill; decltype(value) val; } dill_ = { .val = value }; |
There was a problem hiding this comment.
If we push this union declaration into the if constexpr, then we will only need one integer member with the correct width.
There was a problem hiding this comment.
Doesn't match_any() need to handle int and long?
There was a problem hiding this comment.
Yeah, I was kinda assuming the compiler will optimize out the upper bits since the whole function is a template anyway. But this works too.
* SWDEV-460151 - add gfx1201 to amd-staging clr * SWDEV-460151 - removed pal macro --------- Co-authored-by: Jimbo Xie <[email protected]>
amd-hhashemi
changed the title
| } else { | ||
| union dill { unsigned int i[2]; unsigned long long ill; decltype(value) val; } dill_ = { .val = value }; | ||
| retval = 1; | ||
| //Do a full rotate of the wave lanes, using dpp with "wave_rol1" control (ID: 0x134). |
There was a problem hiding this comment.
| //Do a full rotate of the wave lanes, using dpp with "wave_rol1" control (ID: 0x134). | |
| // Do a full rotate of the wave lanes, using dpp with "wave_rol1" control (ID: 0x134). |
Please add a space after the // for every comment line. I don't know if HIP has yet adopted a coding style, but this kind of whitespace is very commonly considered correct.
| //In doing so each lane gets a mask of matches with all other lanes in the wave in retval. | ||
| for (int i = 1; i < static_cast<int>(warpSize); i++) { | ||
| if constexpr (sizeof(value) == 8) | ||
| dill_.ill = __builtin_amdgcn_mov_dpp(dill_.ill, 0x134 /*dpp_ctrl=wave_rol1*/, 0xf/*row_mask*/, 0xf/*clmn_mask*/, 1/*bound_ctrl*/); |
There was a problem hiding this comment.
The comments are partially helpful, but I am still not okay with magic constants appearing somewhere in code. Those constants like 0x134 need to be declared in one suitable place, maybe as an enum or as constant integers with meaningful names.
|
Imported to ROCm/rocm-systems |
Associated JIRA ticket number/Github issue number
Fixes SWDEV-546808
What type of PR is this? (check all applicable)
What were the changes?
match_any() showing suboptimal perf. This change converts it to a sequence of wave_rot1 dpp ops.
Why are these changes needed?
Prior solution uses inefficient firstlane in loop.
Updated CHANGELOG?
Added/Updated documentation?
Additional Checks